// GPU Verlet Integrator
// Autor: Juan G. Campa

#ifndef _VERLETINTEGRATOR1_
#define _VERLETINTEGRATOR1_

// includes, C string library
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

// includes, cuda
#include <cuda.h>
#include <builtin_types.h>
#include <cuda_runtime_api.h>

// includes, project
//#include <cutil_inline.h>

__device__ float4 operator+(const float4 &a, const float4 &b) {
   return make_float4(a.x+b.x, a.y+b.y, a.z+b.z, a.w+b.w);
}

__device__ float4 operator-(const float4 &a, const float4 &b) {
   return make_float4(a.x-b.x, a.y-b.y, a.z-b.z, a.w-b.w);
}

extern "C" __global__ void Integrate(float4 *posArray, float4 *prevPosArray, const int particleCount, const int dt)
{
   unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;

   /*float4 aux = posArray[i];
   float4 a = aux + aux;
   prevPosArray[i] = aux;*/
   
   if (i < particleCount)
   {
      float4 curPos = posArray[i];
      float4 nextPos = curPos + curPos;
      float4 prevPos = prevPosArray[i];
      nextPos = nextPos - prevPos;
      posArray[i] = nextPos;
      prevPosArray[i] = curPos;
   }
}

extern "C" 
void ExecuteIntegrate(float4* pos, float4* prevPos, unsigned int particleCount, float dt)
{
    cudaError_t error = cudaSuccess;

    dim3 block(256, 1, 1);
    dim3 grid(particleCount / block.x, 1, 1);

	Integrate<<<grid, block>>>(pos, prevPos, particleCount, dt);

    error = cudaGetLastError();
    if (error != cudaSuccess) {
        printf("kernel() failed to launch error = %d\n", error);
    }
}

#endif